#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics: enable
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics: enable
#pragma OPENCL EXTENSION cl_amd_printf : enable


__kernel void
test(__global int* const dData)
{
	const size_t bx = get_group_id(0);
	const size_t tx = get_local_id(0);
	const unsigned index = bx * 256 + tx;
	dData[index] += 1;
}


__kernel void
testA(__global unsigned* const dData, __global volatile unsigned* synch)
{
	const size_t bx = get_group_id(0);
	const size_t tx = get_local_id(0);
	const unsigned index = bx * 256 + tx;
    for (unsigned i = 0; i < 10; ++i) {
          
        //printf("%d\n", *synch);
	    dData[index] += 1;

		if (tx == 0) {
			atom_add(synch, 1);
			while (*synch != (32 * (i + 1))) {
				dData[index] += 1;

			}
		}

		barrier(CLK_GLOBAL_MEM_FENCE);
	}  
}

#if 0
__kernel void
testA(__global unsigned* const dData)
{
    const size_t bx = get_group_id(0);
    const size_t tx = get_local_id(0);
    const unsigned index = bx * 256 + tx;
    for (unsigned i = 0; i < 10; ++i) {

        //printf("%d\n", *synch);
        dData[index] += 1;

        if (tx == 0) {
            atom_add(&synch, 1);
            printf("%d %d %d\n", bx, tx, synch);
            while (synch != 32);
        }

        //barrier(CLK_GLOBAL_MEM_FENCE);
        break;
    }
}

#endif


